#pragma once

#include <cmath>

// Reference:
// https://forums.developer.nvidia.com/t/a-very-fast-low-precision-implementation-of-erff-for-compute-capability-7-5/275503

namespace solar::cuda::fast_math
{
    namespace internal
    {
#if __CUDA_ARCH__ >= 750
        __forceinline__ __device__ float raw_tanhf(float a)
        {
            asm("tanh.approx.f32 %0,%1; \n\t" : "+f"(a) : "f"(a));
            return a;
        }
#endif
    } // namespace internal
#if __CUDA_ARCH__ >= 750
    __forceinline__ __device__ float erff(float x)
    {
        float x2 = x * x;
        x = std::fmaf(std::fmaf(0.100646973f, x2, 0.128759325f), x,
                      x); // 0x1.9c4000p-4, 0x1.07b2f8p-3
        return internal::raw_tanhf(x);
    }
#else
    using std::erff;
#endif
} // namespace solar::cuda::fast_math
